 ///////////////////////////////////////////////////////////////////////////////
//! @file	: face_detect.cl
//! @date   : 2016/03/16
//! @author: guyadong
//! @brief  : detect face in integral matrix
////////////////////////////////////////////////////////////////////////////////
#include "adaboost_face.h"
#include "rect_utilits.h"
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable

/////////////////////////////////////////////////////////////////////////
inline float4 obj_normalization(
	 const int4 obj // xy: obj position,z: detected object sum of obj,w win_size
	,const float2 normalization_factor
	){
	float4 norm_obj=convert_float4(obj);
	norm_obj.xy*=normalization_factor;
	norm_obj.w=(float)(obj.w)*max(normalization_factor.x,normalization_factor.y);
	return norm_obj;
}
inline int4 obj_unnormalization(
	 const float4 norm_obj // xy: obj position,z: detected object sum of obj,w win_size
	,const int2 img_size
	){
	int4 obj=convert_int4(norm_obj);
	obj.xy=convert_int2(norm_obj.xy*convert_float2(img_size));
	obj.w=(int)(norm_obj.w*min(img_size.x,img_size.y));
	return obj;
}

///////////////////////////////////////////////////////////////////////////////
//! @brief  : return true if find intersectant rectangle with norm_obj in detected_objects_buffer_norm 
////////////////////////////////////////////////////////////////////////////////
inline int find_intersect_norm(
	 const float4 norm_obj // xy: obj position,z: object sum of probe,w win_size
	,__global detected_objects_buffer_norm* detected_obj
	,float4 *intersect_obj
		){
	float4 rect1,rect2;
	for(int i=0;i<detected_obj->detected_num;++i){
		rect1=(float4)(norm_obj.xy,norm_obj.xy+norm_obj.w);		
		float4 detected_obj_i=vload4(i,(__global float*)detected_obj->storage);
		rect2=(float4)(detected_obj_i.xy,detected_obj_i.xy+detected_obj_i.w);
		if (is_intersectant_rectangle_norm(rect1,rect2)){			
			*intersect_obj=detected_obj_i;
			return i;
		}		
	}
	return -1;
}

///////////////////////////////////////////////////////////////////////////////
//! @brief  : return true if find intersectant rectangle with obj in detected_objects_buffer_norm 
////////////////////////////////////////////////////////////////////////////////
inline int find_intersect(
	 const int4 obj // xy: obj position,z: object sum of probe,w win_size
	,__global detected_objects_buffer_norm* detected_obj
	,const int2 img_size
	,int4 *intersect_obj
		){	
	const float2 normalization_factor=get_normalization_factor(img_size);
	float4 norm_obj=obj_normalization(obj,normalization_factor);
	float4 intersect_obj_norm;
	int b=find_intersect_norm(norm_obj,detected_obj,&intersect_obj_norm);
	*intersect_obj=obj_unnormalization(intersect_obj_norm,img_size);
	return b;
}


//////////////////////////////////////////////////////////////////////////
inline int4 get_detect_rectangle(int2 pos,int win_size){
	return (int4)(pos,pos+win_size);
}
inline float4 get_detect_rectangle_norm(int2 pos,int win_size,const float2 normalization_factor){
	return rectangle_normalization(get_detect_rectangle(pos,win_size),normalization_factor);
}
inline float4 get_object_rectangle_norm(float4 obj,const float2 normalization_factor){	
	return (float4)(obj.xy,obj.xy+(float2)(obj.w)*normalization_factor/max(normalization_factor.x,normalization_factor.y));
}

typedef enum _rect_intersect_type{
	 NOT_INTERSECT
	,HAS_INTERSECT_OBJECTS
	,INCLUDED_BY_OBJECT
	,INCLUDE_OBJECT
}rect_intersect_type;
///////////////////////////////////////////////////////////////////////////////
//! @brief  : compute intersect status between detect_rect_norm and objectes of detected_obj
//            return rect_intersect_type  
//			  inclued_obj_rect_norm is valid only where intersect_status !=NOT_INTERSECT
////////////////////////////////////////////////////////////////////////////////
inline bool find_intersect_rect_norm(
	 const float4 detect_rect_norm 
	,__constant detected_objects_buffer_norm* detected_obj
	,const float2 normalization_factor
		){
	float4 rect2,detected_obj_i;
	bool res=false;
	for(int i=0;(!res)&&i<detected_obj->detected_num;++i){
		detected_obj_i=vload4(i,(__constant float*)detected_obj->storage);
		rect2=get_object_rectangle_norm(detected_obj_i,normalization_factor);
		res=is_valid_rectangle_norm(intersect_rectangle_norm(detect_rect_norm,rect2));
	}
	return res;
}

///////////////////////////////////////////////////////////////////////////////
//! @brief  : first face filter	
////////////////////////////////////////////////////////////////////////////////
// +-------------------------------------------+
// |                                           |
// |                                           |
// |      A1+-----------------------+A2        |
// |        |                       |          |
// |        |   x1+-----------+x2   |          |
// |        |     |           |     |          |
// |        |     |           |     |          |
// |        |   x3+-----------+x4   |          |
// |        |     |           |     |          |
// |        |     |           |     |          |
// |        |   x5+-----------+x6   |          |
// |        |                       |          |
// |      A3+-----------------------+A4        |
// |                                           |
// |                                           |
// |                                           |
// +-------------------------------------------+
__attribute__((vec_type_hint(float4)))
__kernel void  create_mean( 
 		 matrix_info_cl im_info
 		,const __global  float *integ_mat
 		,matrix_info_cl om_info
 		,__global float * mean_mat
		,int win_size
		,float inv_pixels_num	
		,int sample_step		
 		 ){	
	__local float4 local_buf[MEAN_LOCAL_SIZE];
	const int pixel_size=MEAN_LOCAL_SIZE<<2;
	const int group_y= (int)get_group_id(1);
	const int group_x= (int)get_group_id(0)*pixel_size;
	const int group_count=min(pixel_size,(int)(om_info.width-group_x));
	const int local_id=get_local_id(0);
	float4 mean;
	// A1
	#define INDEX_A1 (group_y*im_info.row_stride+group_x)*sample_step
	//A2
	#define INDEX_A2 (group_y*im_info.row_stride+group_x)*sample_step+win_size
	//A3
	#define INDEX_A3 (group_y*sample_step+win_size)*im_info.row_stride+group_x*sample_step
	//A4
	#define INDEX_A4 (group_y*sample_step+win_size)*im_info.row_stride+group_x*sample_step+win_size
	local_buf[local_id]=0;
	event_t evt;
	barrier(CLK_LOCAL_MEM_FENCE);	
	evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A4,group_count,sample_step,0);wait_group_events(1,&evt);
	mean =local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A1,group_count,sample_step,0);wait_group_events(1,&evt);
	mean+=local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A2,group_count,sample_step,0);wait_group_events(1,&evt);
	mean-=local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A3,group_count,sample_step,0);wait_group_events(1,&evt);
	mean-=local_buf[local_id];
	local_buf[local_id] = mean*inv_pixels_num;
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy(mean_mat+group_y*om_info.row_stride+group_x,(__local float*)local_buf,group_count,0);wait_group_events(1,&evt);
	#undef INDEX_A1
	#undef INDEX_A2
	#undef INDEX_A3
	#undef INDEX_A4
}

__attribute__((vec_type_hint(float4)))
__kernel void create_variance( 
 		matrix_info_cl im_info
 		,const __global  float *intsqr_mat
 		,matrix_info_cl om_info
 		,__global float* mean_mat
 		,__global float* variance_mat		
		,int win_size
		,float inv_pixels_num
		,int sample_step
 		 ){	
	__local float4 local_buf[VARIANCE_LOCAL_SIZE];
	const int pixel_size=MEAN_LOCAL_SIZE<<2;
	const int group_y= (int)get_group_id(1);
	const int group_x= (int)get_group_id(0)*pixel_size;
	const int group_count=min(pixel_size,(int)(om_info.width-group_x));
	const int local_id=get_local_id(0);
	float4 variance;

	// A1
	#define INDEX_A1 (group_y*im_info.row_stride+group_x)*sample_step
	//A2
	#define INDEX_A2 (group_y*im_info.row_stride+group_x)*sample_step+win_size
	//A3
	#define INDEX_A3 (group_y*sample_step+win_size)*im_info.row_stride+group_x*sample_step
	//A4
	#define INDEX_A4 (group_y*sample_step+win_size)*im_info.row_stride+group_x*sample_step+win_size
	local_buf[local_id]=0;	
	event_t evt;
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,intsqr_mat+INDEX_A4,group_count,sample_step,0);wait_group_events(1,&evt);
	variance =local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,intsqr_mat+INDEX_A1,group_count,sample_step,0);wait_group_events(1,&evt);
	variance+=local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,intsqr_mat+INDEX_A2,group_count,sample_step,0);wait_group_events(1,&evt);
	variance-=local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_strided_copy((__local float*)local_buf,intsqr_mat+INDEX_A3,group_count,sample_step,0);wait_group_events(1,&evt);
	variance-=local_buf[local_id];
	variance*=inv_pixels_num;
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy((__local float*)local_buf,mean_mat+group_y*om_info.row_stride+group_x,group_count,0);wait_group_events(1,&evt);
	variance-=pown(local_buf[local_id],2);
	local_buf[local_id]=( variance>(float)(VARIANCE_MINVAR*VARIANCE_MINVAR)
						&&variance<(float)(VARIANCE_MAXVAR*VARIANCE_MAXVAR))
						 ?variance:(float)0;
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy(variance_mat+group_y*om_info.row_stride+group_x,(__local float*)local_buf,group_count,0);wait_group_events(1,&evt);
	#undef INDEX_A1
	#undef INDEX_A2
	#undef INDEX_A3
	#undef INDEX_A4
}
__kernel void create_hit16( 
 		matrix_info_cl om_info
 		,__global uchar16* hit_mat
 		,__global float16* variance_mat
 		 ){
	size_t y = get_global_id(1), x = get_global_id(0);
	float16 variance= variance_mat[(y*om_info.row_stride>>4)+x];
	hit_mat[(y*om_info.row_stride>>4)+x]=convert_uchar16(
											(variance>(float)(VARIANCE_MINVAR*VARIANCE_MINVAR))
										&&  (variance<(float)(VARIANCE_MAXVAR*VARIANCE_MAXVAR))
												);
}
__kernel void intersect_filter( 
 		 __constant detected_objects_buffer_norm* detected_obj_buf_ptr 
		,matrix_info_cl im_info
 		,matrix_info_cl om_info
 		,__global char* hit_mat
		,int win_size		
		,float2 normalization_factor		
		,int sample_step
 		 ){
	const int2 pos= (int2)((int)get_global_id(1),(int)get_global_id(0));
	int4 detect_rect=(int4)(pos*sample_step,pos*sample_step+win_size);
	if(find_intersect_rect_norm(
			rectangle_normalization(detect_rect,normalization_factor)
			,detected_obj_buf_ptr
			,normalization_factor))
		hit_mat[pos.y*om_info.row_stride+pos.x]=0;
	}
	
///////////////////////////////////////////////////////////////////////////////
//! @brief  : compute difference between 'feather' and integral matrix specified by 'detect_win' on 'integral_mat'  	
////////////////////////////////////////////////////////////////////////////////
inline float get_difference(
		 const feather_short_cl feather
		,const matrix_info_cl im_info
 		,const __global  float * const integral_mat
 		,const int2 detect_win
 		){
 	uint tt1,tt2,tt3,tt4,tt5,tt6;
	uint x1,x2,x3,x4,x5,x6,x7,x8,x9;
	//float difference=0;
 	switch(feather.RectangleNum){
	case 2:
		//  +-----------------------+
		//  |                       |
		//  |   x1+-----------+x2   |
		//  |     |           |     |
		//  |     |           |     |
		//  |   x3+-----------+x4   |
		//  |     |           |     |
		//  |     |           |     |
		//  |   x5+-----------+x6   |
		//  |                       |
		//  +-----------------------+
		if(feather.Up_Down){
			tt1 = ((detect_win.y + feather.y)*im_info.row_stride);
			tt2 = ((detect_win.y + feather.CorrectNum)*im_info.row_stride);
			tt3 = ((detect_win.y + feather.NonfaceCorrectNum)*im_info.row_stride);
			tt4 = (detect_win.x + feather.Valve);
			tt5 = (detect_win.x + feather.x);
			x1 = tt1 + tt5;
			x2 = tt1 + tt4;
			x3 = tt5 + tt2;
			x5 = tt5 + tt3;
			x4 = tt4 + tt2;
			x6 = tt4 + tt3;
			/*
			difference+=integral_mat[x6],
			difference-=integral_mat[x5],
			difference-=integral_mat[x4]*2,
			difference+=integral_mat[x3]*2,
			difference+=integral_mat[x2],
			difference-=integral_mat[x1];
			return difference;*/
			return  (integral_mat[x6]-integral_mat[x5])-(integral_mat[x4]-integral_mat[x3])*2  +  (integral_mat[x2]-integral_mat[x1]);
			//return  (x6-x5)-(x4-x3)*2  +  (x2-x1);
		}
		//  +-----------------------+
		//  |    x1     x2    x3    |
		//  |     +-----+-----+     |
		//  |     |     |     |     |
		//  |     |     |     |     |
		//  |     |     |     |     |
		//  |     |     |     |     |
		//  |     |     |     |     |
		//  |     +-----+-----+     |
		//  |    x4     x5    x6    |
		//  +-----------------------+	
		else{
			tt1 = ((detect_win.y + feather.y)*im_info.row_stride);
			tt2 = ((detect_win.y + feather.Valve)*im_info.row_stride);
			tt3 = (detect_win.x + feather.x);
			tt4 = (detect_win.x + feather.CorrectNum);
			tt5 = (detect_win.x + feather.NonfaceCorrectNum);
			x1 = tt1 + tt3;
			x2 = tt1 + tt4;
			x3 = tt1 + tt5;
			x4 = tt2 + tt3;
			x5 = tt2 + tt4;
			x6 = tt2 + tt5;
			/*difference+=integral_mat[x6],
			difference-=integral_mat[x3],
			difference-=integral_mat[x5]*2,
			difference+=integral_mat[x2]*2,
			difference+=integral_mat[x4],
			difference-=integral_mat[x1];
			return difference;*/
			return (integral_mat[x6]-integral_mat[x3]) - ((integral_mat[x5]-integral_mat[x2])*2) + (integral_mat[x4]-integral_mat[x1]);
			//return (x6-x3) - ((x5-x2)*2) + (x4-x1);
		}	
	case 3:
		//  +-----------------------------+
		//  |    x1     x2    x3    x4    |
		//  |     +-----+-----+-----+     |
		//  |     |     |     |     |     |
		//  |     |     |     |     |     |
		//  |     |     |     |     |     |
		//  |     |     |     |     |     |
		//  |     |     |     |     |     |
		//  |     +-----+-----+---- +     |
		//  |    x5     x6    x7    x8    |
		//  +-----------------------------+
		tt1 = ((detect_win.y + feather.y)*im_info.row_stride);
		tt2 = ((detect_win.y + feather.y+feather.Height)*im_info.row_stride);
		tt3 = (detect_win.x + feather.x);
		tt4 = (detect_win.x + feather.CorrectNum);
		tt5 = (detect_win.x + feather.NonfaceCorrectNum);
		tt6 = (detect_win.x + feather.Valve);
		x1 = tt1 + tt3;
		x2 = tt1 + tt4;
		x3 = tt1 + tt5;
		x4 = tt1 + tt6;
		x5 = tt2 + tt3;
		x6 = tt2 + tt4;
		x7 = tt2 + tt5;
		x8 = tt2 + tt6;
		/*difference+=integral_mat[x6],
		difference-=integral_mat[x2],
		difference+=integral_mat[x3],
		difference-=integral_mat[x7],
		difference*=3,
		difference+=integral_mat[x1],
		difference-=integral_mat[x5],
		difference+=integral_mat[x8],
		difference-=integral_mat[x4];
		return difference;*/
		return (integral_mat[x6]-integral_mat[x2]+integral_mat[x3]-integral_mat[x7])*3+integral_mat[x1]-integral_mat[x5]+integral_mat[x8]-integral_mat[x4];
		//return (x6-x2+x3-x7)*3+x1-x5+x8-x4;
	case 4:
		//  +-----------------------+
		//  |    x1     x2    x3    |
		//  |     +-----+-----+     |
		//  |     |     |     |     |
		//  |     |     |     |     |
		//  |   x4+---x5+---x6+     |
		//  |     |     |     |     |
		//  |     |     |     |     |
		//  |     +-----+-----+     |
		//  |    x7     x8    x9    |
		//  +-----------------------+	
		tt1 = ((detect_win.y + feather.y)*im_info.row_stride);
		tt2 = ((detect_win.y + feather.y+(feather.unitHeight))*im_info.row_stride);
		tt3 = ((detect_win.y + feather.y+feather.Height)*im_info.row_stride);
		tt4 = (detect_win.x + feather.x);
		tt5 = (detect_win.x + feather.x+(feather.unitWidth));
		tt6 = (detect_win.x + feather.x+feather.Width);
		x1 = tt1 + tt4;
		x2 = tt1 + tt5;
		x3 = tt1 + tt6;
		x4 = tt2 + tt4;
		x5 = tt2 + tt5;
		x6 = tt2 + tt6;
		x7 = tt3 + tt4;
		x8 = tt3 + tt5;
		x9 = tt3 + tt6;
		/*difference+=integral_mat[x5]*2,
		difference-=integral_mat[x4],
		difference-=integral_mat[x8],
		difference-=integral_mat[x6],
		difference-=integral_mat[x2],
		difference*=2,
		difference+=integral_mat[x1],
		difference+=integral_mat[x9],
		difference+=integral_mat[x7],
		difference+=integral_mat[x3];
		return difference;*/
		return 
			((((integral_mat[x5]*2-integral_mat[x4]-integral_mat[x8]-integral_mat[x6]-integral_mat[x2])*2) + integral_mat[x1] + integral_mat[x9])
			+
			(integral_mat[x7] +integral_mat[x3]));
		//return 
		//	((((x5*2-x4-x8-x6-x2)*2) + x1 + x9)
		//	+
		//	(x7 +x3));
	default:
		return (float)(0);
	} 	
 }
///////////////////////////////////////////////////////////////////////////////
//! @brief  : compute difference between 'feather' and 'variance' ,'difference'	
////////////////////////////////////////////////////////////////////////////////
inline float compute_correct(
		 const feather_short_cl feather
 		,const float weight   // weight of 'feather'
		,const float variance  
		,const float difference
 		){ 		
	if(feather.Sign){
		if (feather.DValve > 0){
			return (difference > 0.0f)&&((pown(feather.DValve,2) * variance) < pown(difference, 2) ) ? weight:(float)(0.0);
		}else{
			return  ((difference <= 0)&&(pown(feather.DValve,2) * variance > pown(difference, 2) ))||(difference > 0)? weight:(float)(0.0);
		}
	}else	{
		if (feather.DValve > 0){
			return  ((difference > 0)&&(pown(feather.DValve,2) * variance > pown(difference, 2)))||(difference <= 0)? weight:(float)(0.0);
		}else	{
			return (difference <= 0)&&(pown(feather.DValve,2) * variance < pown(difference, 2))? weight:(float)(0.0);
		}
	}	
}

__attribute__((vec_type_hint(float)))
__attribute__((reqd_work_group_size(HAAR_FILTER_LOCAL_WIDTH, HAAR_FILTER_LOCAL_HEIGHT, 1)))
__kernel void haar_filter( 
  		__constant casinfo_base_float* const cas
		,__constant feather_short_cl* const_feathers
		,feather_block_base block
 		,matrix_info_cl im_info
 		,const __global  float *integ_mat
 		,matrix_info_cl om_info
 		,__global char* hit_mat
		,__global float* variance_mat
		,__global char* hit_mat_out
		,int sample_step	
 		 ){
	__local int2 group_pos	[HAAR_FILTER_LOCAL_WIDTH*HAAR_FILTER_LOCAL_HEIGHT];
	__local int2 detect_win	[HAAR_FILTER_LOCAL_WIDTH*HAAR_FILTER_LOCAL_HEIGHT];
	__local int2 detect_win_group;
	__local int hit_count;
	feather_short_cl feather_obj;
	float correct;	
	int local_index=get_local_id(1)*HAAR_FILTER_LOCAL_WIDTH+get_local_id(0);
	group_pos[local_index]=(int2)((int)get_local_id(0),(int)get_local_id(1));
	detect_win_group=(int2)((int)(get_group_id(0)*HAAR_FILTER_LOCAL_WIDTH),(int)(get_group_id(1)*HAAR_FILTER_LOCAL_HEIGHT));
	hit_count=0;
	detect_win_group*=sample_step;
	barrier(CLK_LOCAL_MEM_FENCE);
	detect_win[local_index]=detect_win_group+group_pos[local_index]*sample_step;
	bool in_range= (int)get_local_id(0)<min((int)HAAR_FILTER_LOCAL_WIDTH ,(int)om_info.width -(int)get_group_id(0)*HAAR_FILTER_LOCAL_WIDTH)
				 &&(int)get_local_id(1)<min((int)HAAR_FILTER_LOCAL_HEIGHT,(int)om_info.height-(int)get_group_id(1)*HAAR_FILTER_LOCAL_HEIGHT);

	char hit=in_range?hit_mat[get_global_id(1)*om_info.row_stride + get_global_id(0)]:0;
	if(hit)
		atom_inc(&hit_count);
	barrier(CLK_LOCAL_MEM_FENCE);
	for(int layer=block.start_layer,feather_index=block.offset,end_layer=block.end_layer;hit_count&&layer<end_layer ;feather_index += cas->FeaNumPerLayer[layer++]){		
		correct = (float)(0);
		for( int sub_fea=feather_index,end_sub=feather_index+cas->FeaNumPerLayer[layer];sub_fea < end_sub;++sub_fea){
			feather_obj=const_feathers[sub_fea];
			barrier(CLK_LOCAL_MEM_FENCE);			
			float difference = hit?get_difference( 
				feather_obj
				,im_info
 				,integ_mat
 				,detect_win[local_index]
 				):0;
			correct += hit?compute_correct( 
				feather_obj
				,cas->Weight[sub_fea]
				,variance_mat[get_global_id(1)*om_info.row_stride + get_global_id(0)]
				,difference
				):0;
		}
		if(hit&&correct < cas->LayerWeight[layer]){
			hit=0;
			atom_dec(&hit_count);
		}
		barrier(CLK_LOCAL_MEM_FENCE);
	}
	if(hit)	
		hit_mat_out[get_global_id(1)*om_info.row_stride + get_global_id(0)]=0xff;
}

__kernel void create_hit_mat_full( 
 		 matrix_info_cl om_info
 		,__global uchar* hit_mat
		,matrix_info_cl om_info_full
 		,__global uchar* hit_mat_full
		,int sample_step
 		 ){
	int2 pos_src=(int2)((int)get_global_id(0),(int)get_global_id(1));
	int2 pos_dst=pos_src*sample_step;
	if(hit_mat[pos_src.y*om_info.row_stride+pos_src.x])
		hit_mat_full[pos_dst.y*om_info_full.row_stride+pos_dst.x]=0xff;
}

///////////////////////////////////////////////////////////////////////////////
//! @brief  : compute face density on integral count matrix created by hit matrix	
////////////////////////////////////////////////////////////////////////////////
// +-------------------------------------------+
// |                                           |
// |                                           |
// |      A1+-----------------------+A2        |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |        |                       |          |
// |      A3+-----------------------+A4        |
// |                                           |
// |                                           |
// |                                           |
// +-------------------------------------------+
__attribute__((vec_type_hint(ushort8)))
__kernel void object_density_filter( 
 		  matrix_info_cl im_info
 		, __global  ushort *integ_count_mat
 		, matrix_info_cl om_info
 		, __global ushort8* density_mat
		,int  face_dist_size
		,ushort sum_threshold
 		 ){
	__local ushort8 local_buf[DENSITY_FILTER_LOCAL_SIZE];
	int group_index=(int)(get_group_id(0)*DENSITY_FILTER_LOCAL_SIZE);		 
	int group_count=min((int)DENSITY_FILTER_LOCAL_SIZE,(int)(((om_info.width+7)>>3)-group_index));
	int local_id=get_local_id(0);
	const int y	= (int)get_global_id(1);
	int index_a1=y*im_info.row_stride+(group_index<<3);						//A1
	int index_a2=index_a1+face_dist_size;									//A2
	int index_a3=index_a1+face_dist_size*im_info.row_stride;	            //A3
	int index_a4=index_a3+face_dist_size;									//A4
	ushort8 sum;
	event_t evt;
	// compute faces sum in each window which size speciialed win_size : A4+A1-A2-A3
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy((__local ushort *)local_buf,integ_count_mat+index_a4,group_count<<3,0);wait_group_events(1,&evt);
	sum =local_buf[local_id];
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy((__local ushort *)local_buf,integ_count_mat+index_a1,group_count<<3,0);wait_group_events(1,&evt);
	sum+=local_buf[local_id],
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy((__local ushort *)local_buf,integ_count_mat+index_a2,group_count<<3,0);wait_group_events(1,&evt);
	sum-=local_buf[local_id],
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy((__local ushort *)local_buf,integ_count_mat+index_a3,group_count<<3,0);wait_group_events(1,&evt);
	sum-=local_buf[local_id];
	local_buf[local_id]=sum>=sum_threshold?sum:(ushort8)(0);
	barrier(CLK_LOCAL_MEM_FENCE);
	evt=async_work_group_copy(density_mat+(y*om_info.row_stride>>3)+group_index,local_buf,group_count,0);wait_group_events(1,&evt);
}

__attribute__((reqd_work_group_size(8, 8, 1)))
__kernel void object_cluster_group(
 		  matrix_info_cl density_mat_info
 		, const __global ushort* density_mat
		, __global detected_objects_buffer_level* buffer_level
		, __constant cluster_filter_const_param1 * param
 		 ){
	// x,y object location (normalized coordinate)
	// z, object sum
	__local int max_density;
	max_density=0;
	barrier(CLK_LOCAL_MEM_FENCE);
	ushort2 pos=     (ushort2)((ushort)get_global_id(0),(ushort)get_global_id(1))*(ushort)param->sample_step;
	short2 group_pos=( short2)(( short)get_group_id (0),( short)get_group_id (1))<<3;
	short2 group_count=min((short2)(8)
						  ,(short2)((short)CEIL_DIV(density_mat_info.width ,param->sample_step)-group_pos.x
						           ,(short)CEIL_DIV(density_mat_info.height,param->sample_step)-group_pos.y));
	short2 local_pos=(short2)((short)get_local_id(0),(short)get_local_id(1));
	ushort4 pos_density=all(group_count>(short)0)&&all(local_pos<group_count)?(ushort4)(pos+(ushort)((param->face_dist_size>>1)+1),(ushort)density_mat[pos.y*density_mat_info.row_stride+pos.x],(ushort)0):(ushort)0;
	if(pos_density.z>0){
		atom_max(&max_density,pos_density.z);
	}
	barrier(CLK_LOCAL_MEM_FENCE);
	if(max_density&&pos_density.z==max_density){
		int c=atom_inc(&buffer_level->detected_num);
		if(c<MAX_DETECTED_OBJECT_NUM_LEVEL)
			buffer_level->storage[c]=pos_density;
		else
			buffer_level->status=KERNEL_OUT_INDEX;		
	}
}
__kernel void object_cluster(
		  __constant detected_objects_buffer_level* buffer_level
		, __global detected_objects_buffer_norm* buffer_norm
		, __constant cluster_filter_const_param1 * param
 		 ){
	const int gid=get_global_id(0);
	if(gid>=buffer_level->detected_num)return;
	const int win_size=param->win_size;
	int4 obj=convert_int4(buffer_level->storage[gid]);	
	int4 obj_rect=(int4)(obj.xy,obj.xy+win_size);
	const int obj_len=obj_rect.x*obj_rect.x+obj_rect.y*obj_rect.y;	
	for(int i=0,end_i=buffer_level->detected_num;i<end_i;++i){
		if(gid==i)continue;
		int4 other=convert_int4(buffer_level->storage[i]);		
		int4 other_rect=(int4)(other.xy,other.xy+win_size);
		if(is_intersectant_rectangle(obj_rect,other_rect)){
			if(obj.z<other.z)return;
			else if(obj.z==other.z&&obj_len<other_rect.x*other_rect.x+other_rect.y*other_rect.y)return;
		}
	}
	int c=atom_inc(&buffer_norm->detected_num);
	if(c<MAX_DETECTED_OBJECT_NUM){
		obj.w=win_size;
		float4 obj_norm=obj_normalization(obj,param->normalization_factor);
		buffer_norm->storage[c]=obj_norm;
	}else
		buffer_norm->status=KERNEL_OUT_INDEX;
}
